
char clCode_clppSort_BitonicSort[]=
"#define WGZ 32\n"
"#define WGZ_x2 (WGZ*2)\n"
"#define WGZ_x3 (WGZ*3)\n"
"#define WGZ_x4 (WGZ*4)\n"
"#define WGZ_1 (WGZ-1)\n"
"#define WGZ_2 (WGZ-2)\n"
"#define WGZ_x2_1 (WGZ_x2-1)\n"
"#define WGZ_x3_1 (WGZ_x3-1)\n"
"#define WGZ_x4_1 (WGZ_x4-1)\n"
"#ifdef KEYS_ONLY\n"
"#define KEY(DATA) (DATA)\n"
"#else\n"
"#define KEY(DATA) (DATA.s0)\n"
"#endif\n"
"#define BARRIER_LOCAL barrier(CLK_LOCAL_MEM_FENCE)\n"
"__kernel\n"
"void kernel__BitonicSort(__global KV_TYPE* taskIndices, const uint stage, const uint passOfStage)\n"
"{\n"
"	uint sortIncreasing = 1; // Direction\n"
"	uint gid = get_global_id(0);\n"
"	uint pairDistance = 1 << (stage - passOfStage);\n"
"	uint blockWidth = 2 * pairDistance;\n"
"	uint leftId = (gid % pairDistance) + (gid / pairDistance) * blockWidth;\n"
"	uint rightId = leftId + pairDistance;\n"
"	KV_TYPE leftElement = taskIndices[leftId];\n"
"	KV_TYPE rightElement = taskIndices[rightId];\n"
"	\n"
"	uint sameDirectionBlockWidth = 1 << stage;\n"
"	\n"
"	sortIncreasing = ((gid/sameDirectionBlockWidth) % 2 == 1) ? (1 - sortIncreasing) : sortIncreasing;\n"
"	uint leftKey = KEY(taskIndices[leftId]);\n"
"	uint rightKey = KEY(taskIndices[rightId]);\n"
"	\n"
"	KV_TYPE greater = leftKey > rightKey ? leftElement : rightElement;\n"
"KV_TYPE lesser = leftKey > rightKey ? rightElement : leftElement;\n"
"	\n"
"	taskIndices[leftId] = sortIncreasing ? lesser : greater;    \n"
"taskIndices[rightId] = sortIncreasing ? greater : lesser;\n"
"}\n"
;
